//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : prec.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef PREC_H
#define	PREC_H

#include "kernels/prec_kernels.h"

namespace rrbsolver
{

template <class T>
inline void prec_phase1(Grid grid)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texCC, grid.cc, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture CC @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texNN, grid.nn, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)),
        "rrbsolver::bind texture NN @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texEE, grid.ee, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture EE @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texSS, grid.ss, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SS @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texWW, grid.ww, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture WW @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texNE, grid.ne, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture NE @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texSE, grid.se, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SE @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texSW, grid.sw, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SW @ kernel_prec1");
    checkCudaError(cudaBindTexture2D(0, &texNW, grid.nw, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture NW @ kernel_prec1");
#endif

    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_prec1<T> <<< dimGrid, dimBlock >>> (grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_prec1");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texCC), 
        "rrbsolver::unbind texture CC @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texNN), 
        "rrbsolver::unbind texture NN @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texEE), 
        "rrbsolver::unbind texture EE @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texSS), 
        "rrbsolver::unbind texture SS @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texWW), 
        "rrbsolver::unbind texture WW @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texNE), 
        "rrbsolver::unbind texture NE @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texSE), 
        "rrbsolver::unbind texture SE @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texSW), 
        "rrbsolver::unbind texture SW @ kernel_prec1");
    checkCudaError(cudaUnbindTexture(texNW), 
        "rrbsolver::unbind texture NW @ kernel_prec1");
#endif
}

template <class T>
inline void prec_phase2(Grid grid)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texCC, grid.cc, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture CC @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texNN, grid.nn, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)),
        "rrbsolver::bind texture NN @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texEE, grid.ee, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture EE @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texSS, grid.ss, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SS @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texWW, grid.ww, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture WW @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texNE, grid.ne, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture NE @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texSE, grid.se, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SE @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texSW, grid.sw, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SW @ kernel_prec2");
    checkCudaError(cudaBindTexture2D(0, &texNW, grid.nw, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture NW @ kernel_prec2");
#endif

    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_prec2<T> <<< dimGrid, dimBlock >>> (grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_prec2");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texCC), 
        "rrbsolver::unbind texture CC @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texNN), 
        "rrbsolver::unbind texture NN @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texEE), 
        "rrbsolver::unbind texture EE @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texSS), 
        "rrbsolver::unbind texture SS @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texWW), 
        "rrbsolver::unbind texture WW @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texNE), 
        "rrbsolver::unbind texture NE @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texSE), 
        "rrbsolver::unbind texture SE @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texSW), 
        "rrbsolver::unbind texture SW @ kernel_prec2");
    checkCudaError(cudaUnbindTexture(texNW), 
        "rrbsolver::unbind texture NW @ kernel_prec2");
#endif
}

template <class T>
inline void prec_phase3(Grid grid)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texCC, grid.cc, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture CC @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texNN, grid.nn, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)),
        "rrbsolver::bind texture NN @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texEE, grid.ee, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture EE @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texSS, grid.ss, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SS @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texWW, grid.ww, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture WW @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texNE, grid.ne, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture NE @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texSE, grid.se, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SE @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texSW, grid.sw, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture SW @ kernel_prec3");
    checkCudaError(cudaBindTexture2D(0, &texNW, grid.nw, &desc, 
        grid.nx, grid.ny, grid.nx * sizeof(T)), 
        "rrbsolver::bind texture NW @ kernel_prec3");
#endif

    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_prec3<T> <<< dimGrid, dimBlock >>> (grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_prec3");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texCC), 
        "rrbsolver::unbind texture CC @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texNN), 
        "rrbsolver::unbind texture NN @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texEE), 
        "rrbsolver::unbind texture EE @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texSS), 
        "rrbsolver::unbind texture SS @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texWW), 
        "rrbsolver::unbind texture WW @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texNE), 
        "rrbsolver::unbind texture NE @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texSE), 
        "rrbsolver::unbind texture SE @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texSW), 
        "rrbsolver::unbind texture SW @ kernel_prec3");
    checkCudaError(cudaUnbindTexture(texNW), 
        "rrbsolver::unbind texture NW @ kernel_prec3");
#endif
}

template <class T>
inline void prec_phase4(Grid grid)
{
    dimGrid.x = grid.cx / NUM_THREADS_X;
    dimGrid.y = grid.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_prec4<T> <<< dimGrid, dimBlock >>> (grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_prec4");
}

template <class T>
inline void prec_phase5(Grid orig)
{
#if USE_TEXTURES
    checkCudaError(cudaBindTexture2D(0, &texCC, orig.cc, &desc, 
        orig.nx, orig.ny, orig.nx * sizeof(T)), 
        "rrbsolver::bind texture CC @ kernel_prec5");
#endif

    dimGrid.x = orig.cx / NUM_THREADS_X;
    dimGrid.y = orig.cy / NUM_THREADS_Y;
    dimBlock.x = NUM_THREADS_X;
    dimBlock.y = NUM_THREADS_Y;

    kernel_prec5<T> <<< dimGrid, dimBlock >>> (orig);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_prec5");

#if USE_TEXTURES
    checkCudaError(cudaUnbindTexture(texCC), 
        "rrbsolver::unbind texture CC @ kernel_prec5");
#endif
}

template <class T>
inline void prec_final(Grid grid)
{
    dimGrid.x = 1;
    dimGrid.y = 1;
    dimBlock.x = DIM_COMPUTE_BLOCK;
    dimBlock.y = DIM_COMPUTE_BLOCK / 2;

    kernel_precfinal<T> <<< dimGrid, dimBlock >>> (grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_precfinal");
}

} // namespace rrbsolver

#endif // PREC_H

